#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <float.h>
#include <iostream>
#include <stdio.h>



#define Check_Cuda_Runtime(op) __check_cuda_runtime((op),#op,__FILE__,__LINE__)
void __check_cuda_runtime(cudaError_t code,const char* op,const char* file,const char* line){
    if(code!= cudaSuccess){
        const char* cuda_error_name = cudaGetErrorName(code);
        const char* cuda_error_string = cudaGetErrorString(code);
        printf("cuda error happend at file:%s line:%d cudaop %s cudaerroename:%s cudaerroestring:%s",file,line,op,cuda_error_name,cuda_error_string);
    }
}


template<typename T>
__global__ void cuda_histogram(T* A,T* B,int total_elements,bool use_vector){
    int idx = blockIdx.x*blockDim.x + threadIdx.x;
    if(use_vector){
        idx *= 4;
        if(idx<total_elements){
            int4 vec_a = *reintetpret<int4*>(&A[idx]);
            atomicAdd(&B[vec_a.x],1);
            atomicAdd(&B[vec_a.y],1);
            atomicAdd(&B[vec_a.z],1);
            atomicAdd(&B[vec_a.w],1);
        }
    }else{
        if(idx<total_elements){
            atomicAdd(&B[A[idx]],1);
        }
    }
}